#include "matrix3D.cuh"
#include "timer.cuh"
#include <cstdio>

__device__ inline unsigned int lcg_rand(unsigned int& seed)
{
    seed = seed * 1664525 + 1013904223UL;
    return seed;
}

__device__ inline unsigned int taus_rand_step(unsigned int& state,unsigned int S1,unsigned int S2,unsigned int S3,unsigned int M)
{
    unsigned int b = (((state << S1) ^ state) >> S2);
    state = (((state & M) << S3) ^ b);
    return state;
}

__device__ unsigned int taus_rand(uint4& state)
{
    return
        taus_rand_step(state.x, 6U, 13U, 18U, 4294967294UL) ^
        taus_rand_step(state.y, 2U, 27U, 2U, 4294967288UL) ^
        taus_rand_step(state.z, 13U, 21U, 7U, 4294967280UL) ^
        taus_rand_step(state.w, 3U, 12U, 13U, 4294967268UL);

}

__global__ void Hybrid(short * input)
{
    const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
    uint4 sr;
    sr.x= (tid+1) >> (MT_SHIFT0);
    sr.y= (tid+3) << (MT_SHIFTB);
    sr.z= (tid+5) << (MT_SHIFTC);
    sr.w= (tid+7) >> (MT_SHIFT1);
    input[tid]=(short)taus_rand(sr);
}

void controlFunction(short * input,int size)
{
    short * output;
    cudaMalloc((void**)&output,size*2);
    Timer timer;
    cudaEvent_t start = timer.go();
    if((size % 64)==0)
        Hybrid<<<size/512,512>>>(output);
    else
        Hybrid<<<(int)(size/512)+1,512>>>(output);
    printf("Material generation ends with %f Ms \n", timer.get(start));
    cudaMemcpy(input,output,size*2,cudaMemcpyDeviceToHost);
    cudaFree(output);
}
